/*
 *
 *	Yuri Toledo	8124361
 *	Plinio Brito	7696243
 * 	Alexandre
 */

#include <stdio.h>
//APAGAR
/*#define mask_width 5						// Tamanho da mascara (janela) do filtro a ser aplicado
#define mask_radius mask_width/2			// Raio da mascara para controle de loop
#define TILE_WIDTH 16						// Dimensão do bloco de threads (16x16 threads)
#define win (TILE_WIDTH + mask_width - 1)	// Dimensão da janela de dados (window) a ser utilizada por cada bloco
*/

int mask_size = 5; //Tamanho da mascara do filtro
float mask_radius = mask_size/2; //Raio da mascara do filtro
int t_block = 16;	//Blocos de threads
int data_win = (t_block + mask_size - 1); //janela de dados


typedef int pixel;							// Pixel da imagem

int height,width;							// Altura e largura da imagem

pixel *compR;								// Componente vermelha da imagem
pixel *compG;								// Componente verde da imagem
pixel *compB;								// Componente azul da imagem

pixel *inR, *inG, *inB;			// Ponteiros para a imagem de entrada na GPU por canal
pixel *outR, *outG, *outB;		// Ponteiros para a imagem de saida na GPU por canal

 // Kernel para aplicação do filtro. R, G e B são calculadas simultaneamente
__global__ void smoothKernel(const pixel *R, const pixel *G, const pixel *B,
                            pixel *outR, pixel *outG, pixel *outB, 
							const int width, const int height) {
	
	//Memoria compartilhada para janelas de valores utilizados pelo bloco
	__shared__ pixel sR[data_win][data_win]; // RED
	__shared__ pixel sG[data_win][data_win]; // GREEN
	__shared__ pixel sB[data_win][data_win]; // BLUE
	
	int x, y;
	
	// Posicoes para as janelas do bloco e da memoria compartilhada
	int sharedPos, sharedRow,sharedCol;
	int imageRow,imageCol,imagePos;
	
	// loop controla o preenchimento da janela de memoria compartilhada com as threads disponiveis
	int i,loop = (data_win*data_win)/(t_block*t_block)+1;

	for(i=0;i<loop;i++){

		// (data_win*data_win - (t_block*t_block)*i ) = numero de valores que faltam
		
		// Posicao na memoria compartilhada do pixel a ser transferido
		sharedPos = threadIdx.y*t_block + threadIdx.x + (t_block*t_block)*i;
		sharedRow = sharedPos/data_win;
		sharedCol = sharedPos%data_win;
		
		// Posicao do pixel da imagem que sera transiferido
		imageRow = blockIdx.y * t_block + sharedRow - mask_radius;
		imageCol = blockIdx.x * t_block + sharedCol - mask_radius;
		imagePos = imageRow*width + imageCol;
		
		// Verifica se a theread se encontra no intervalo de transferencia
		if (sharedRow < data_win) {
			
			//Pixel é transferido se a posicao for dentro da imagem, caso esteja fora(mascara sendoa aplicada dos pixel da borda).
			// é atribuido o valor zedo para o calcuo
			if (imageRow >= 0 && imageRow < height && imageCol >= 0 && imageCol < width){
				sR[sharedRow][sharedCol] = R[imagePos];
				sG[sharedRow][sharedCol] = G[imagePos];
				sB[sharedRow][sharedCol] = B[imagePos];
			}
			else{
				sR[sharedRow][sharedCol] = 0;
				sG[sharedRow][sharedCol] = 0;
				sB[sharedRow][sharedCol] = 0;
			}
		}
	
	}

    // Sincroniza as threads.
    syncthreads();
 
	// Soma os componentes da mascara
    int somaR=0,somaG=0,somaB=0;
    
	// Aplica a mascara nas componentes R, G e B
    for (y = 0; y < mask_size; y++){
		for (x = 0; x < mask_size; x++){
			somaR += sR[threadIdx.y + y][threadIdx.x + x];
			somaG += sG[threadIdx.y + y][threadIdx.x + x];
			somaB += sB[threadIdx.y + y][threadIdx.x + x];
		}
	}

	// Atribui a media encontrada aos respectivos pixels de cada thread
	y = blockIdx.y * t_block + threadIdx.y;
	x = blockIdx.x * t_block + threadIdx.x;
    if (y < height && x < width){
		outR[y * width + x] = (pixel)somaR/(mask_size*mask_size);
		outG[y * width + x] = (pixel)somaG/(mask_size*mask_size);	
		outB[y * width + x] = (pixel)somaB/(mask_size*mask_size);
	}

    syncthreads();	// Sincroniza threads
}

// Funcao transfere os dados da imagem de entrada para a GPU 
int imageToGPU(){
	
	int error, N = height*width;
	
	// Aloca o espaço na placa de video
	error = cudaMalloc( (void**)&inR, N*sizeof(pixel));
    error |= cudaMalloc( (void**)&inG, N*sizeof(pixel));
    error |= cudaMalloc( (void**)&inB, N*sizeof(pixel));

	error |= cudaMalloc( (void**)&outR, N*sizeof(pixel));
    error |= cudaMalloc( (void**)&outG, N*sizeof(pixel));
    error |= cudaMalloc( (void**)&outB, N*sizeof(pixel));
	
	if(error)
		return 0;

	// Copia imagem para a memoria alocada
	cudaMemcpy(inR, compR, N*sizeof(pixel), cudaMemcpyHostToDevice);
    cudaMemcpy(inG, compG, N*sizeof(pixel), cudaMemcpyHostToDevice);
	cudaMemcpy(inB, compB, N*sizeof(pixel), cudaMemcpyHostToDevice);
	
	return 1;
}

// Volta o resultado do filto para a CPU
void imageToCPU(){
	int N = height*width;

	// Copia o resultado para a CPU
	cudaMemcpy(compR, outR, N*sizeof(pixel), cudaMemcpyDeviceToHost);
    cudaMemcpy(compG, outG, N*sizeof(pixel), cudaMemcpyDeviceToHost);
	cudaMemcpy(compB, outB, N*sizeof(pixel), cudaMemcpyDeviceToHost);

	// Libera memoria alocada na GPU
	cudaFree(outR);
	cudaFree(outG);
	cudaFree(outB);
	cudaFree(inR);
	cudaFree(inG);
	cudaFree(inB);
}

// Aplica o filtro da media de maneira paralela na GPU
void smoothParalelo(){
	int d1 = width/t_block;
	if(width%t_block>0)	// Arredondamento caso nao seja divisivel
		d1++;

	int d2 = height/t_block;
	if(height%t_block>0) // Arredondamento caso nao seja divisivel
		d2++;

	dim3 dimGrid(d1, d2);
	dim3 dimBlock(t_block, t_block);
	cudaError_t cudaStatus;

	if(imageToGPU()){
	
		// Executa kernel CUDA se a alocacao de memoria nao indicar erros
		smoothKernel<<<dimGrid, dimBlock>>>(inR,inG,inB,outR,outG,outB,width,height);
		
		// Verifica se houve erro no kernel
		cudaStatus = cudaGetLastError();
		if (cudaStatus != cudaSuccess)
			fprintf(stderr, "addKernel launch failed: %s\n", cudaGetErrorString(cudaStatus));
		
		// Retorna a imagem resultante para a CPU
		imageToCPU();
	}
	else
		printf("Memória insuficiente. \n");
}

// Aplica o filtro de maneira sequencial
void smoothSeq(){

	int i,j,pos;
	int k,l,somaR,somaG,somaB;
	int r = (int)mask_size/2;
	int div = mask_size*mask_size;

	pixel *outR = (pixel*)malloc(sizeof(pixel)*height*width);
	pixel *outG = (pixel*)malloc(sizeof(pixel)*height*width);
	pixel *outB = (pixel*)malloc(sizeof(pixel)*height*width);
	
	for(i=0;i<height;i++){
		for(j=0;j<width;j++){
			pos = width*i+j;
			
			somaR = 0;
			somaG = 0;
			somaB = 0;
			for(k=i-r;k<=i+r;k++){
				for(l=j-r;l<=j+r;l++){
					if(k>-1 && k<height && l>-1 && l<width){
						somaR+=compR[width*k+l];
						somaG+=compG[width*k+l];
						somaB+=compB[width*k+l];
					}
				}
			}

			outR[pos] = (int)somaR/div;
			outG[pos] = (int)somaG/div;
			outB[pos] = (int)somaB/div;
		}
	}

	compR = outR;
	compG = outG;
	compB = outB;
}

int main(int argc, char** argv){

	char line[200];
	FILE *in,*out;
	int i;

	in = fopen("in.ppm","r");		// Imagem de entrada
	out = fopen("out.ppm","w");		// Imagem de saida

	if(in==NULL){
		printf("Impossivel abrir o arquivo in.ppm. \n");
		return 1;
	}
	if(out==NULL){
		printf("Impossivel abrir o arquivo out.ppm. \n");
		return 1;
	}
	
	fscanf(in,"%s",line); 				// Elimina descritor de formato (P3)

	fscanf(in,"%d%d",&width,&height); 	// le altura e largura em pixels da imagem
	fscanf(in,"%d",&i); 				// elimina descritor de alcance das cores (valor maximo de um pixel)
	
	// Aloca imagem na memoria da CPU 
	compR = (pixel*)malloc(sizeof(pixel)*width*height);
	compG = (pixel*)malloc(sizeof(pixel)*width*height);
	compB = (pixel*)malloc(sizeof(pixel)*width*height);
	
	// Le pixels da imagem
	for(i=0;i<width*height;i++)
		fscanf(in,"%d%d%d",&compR[i],&compG[i],&compB[i]);
	
	// Smooth paralelo
	smoothParalelo();

	// Escreve a imagem resultante
	fprintf(out,"P3\n");
	fprintf(out,"%d %d\n",width,height);
	fprintf(out,"255\n");
	
	for(i=0;i<width*height;i++)
		fprintf(out,"%d\n%d\n%d\n",compR[i],compG[i],compB[i]);
	
	// Finaliza o programa
	fclose(out);
	fclose(in);

	// Libera memoria alocada
	free(compR);
	free(compG);
	free(compB);
	
	return 0;
}
